void OneIter(__global float * pMatrix, __global float * pVector, __global float * pResult, int size, __global float * pNextSolution)
{
	uint row = get_global_id(0);
	uint idMatrix = size * row;
	
	float sum = 0.0f;
	float alpha = 0.0f;
	
	for(uint i=0; i<size; i++)
	{
		sum += pMatrix[idMatrix + i] * pResult[i];
	}
	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
	alpha = 1.0f / pMatrix[idMatrix + row];
	pNextSolution[row] = pResult[row] + alpha * (pVector[row] - sum);
	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
}

float Norm(__global float * pResult, __global float * pNextSolution, int size)
{
	float norm = 0.0f;
	for(uint i=0; i<size; i++)
	{
		float diff = pResult[i] - pNextSolution[i];
		norm += diff * diff;
	}
	barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
	
	norm = sqrt(norm);
	
	return norm;
}
// -----------------------------------------------
// Kernel function
// -----------------------------------------------
__kernel void Main(__global float * pMatrix, __global float * pVector, __global float * pResult, int size, float solvePrecision, __global float * pNextSolution)
{
	float norm = 0.0f;
	for(int i=0; i<111; i++)
	//do
	{
		OneIter(pMatrix, pVector, pResult, size, pNextSolution);
		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
		
		norm = 0.0f;
		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
		
		//uint row = get_global_id(0);
		//float diff = pNextSolution[row] - pResult[row];
		//norm += diff * diff;
		//barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
		
		//norm = sqrt(norm);
		//barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
		
		norm = Norm(pResult, pNextSolution, size);
		
		uint col = get_global_id(0);
		pResult[col] = pNextSolution[col];
		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);
	}
	//while (norm > solvePrecision);
}